<!-- HTML header for doxygen 1.8.13-->
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.14"/>
<meta name="viewport" content="width=device-width, initial-scale=1"/>
<title>Taskflow Handbook</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<link rel="icon" type="image/x-icon" href="favicon.ico" />
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="navtree.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="resize.js"></script>
<script type="text/javascript" src="navtreedata.js"></script>
<script type="text/javascript" src="navtree.js"></script>
<script type="text/javascript">
/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
  $(document).ready(initResizable);
/* @license-end */</script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td id="projectalign" style="padding-left: 0.5em;">
   <div id="projectname"><a href="https://taskflow.github.io/">Taskflow</a>
   &#160;<span id="projectnumber">3.0.0-Master-Branch</span>
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.14 -->
<script type="text/javascript">
/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
var searchBox = new SearchBox("searchBox", "search",false,'Search');
/* @license-end */
</script>
<script type="text/javascript" src="menudata.js"></script>
<script type="text/javascript" src="menu.js"></script>
<script type="text/javascript">
/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
$(function() {
  initMenu('',true,false,'search.php','Search');
  $(document).ready(function() { init_search(); });
});
/* @license-end */</script>
<div id="main-nav"></div>
</div><!-- top -->
<div id="side-nav" class="ui-resizable side-nav-resizable">
  <div id="nav-tree">
    <div id="nav-tree-contents">
      <div id="nav-sync" class="sync"></div>
    </div>
  </div>
  <div id="splitbar" style="-moz-user-select:none;" 
       class="ui-resizable-handle">
  </div>
</div>
<script type="text/javascript">
/* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&amp;dn=gpl-2.0.txt GPL-v2 */
$(document).ready(function(){initNavTree('cuda__capturer_8hpp_source.html','');});
/* @license-end */
</script>
<div id="doc-content">
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div class="header">
  <div class="headertitle">
<div class="title">cuda_capturer.hpp</div>  </div>
</div><!--header-->
<div class="contents">
<a href="cuda__capturer_8hpp.html">Go to the documentation of this file.</a><div class="fragment"><div class="line"><a name="l00001"></a><span class="lineno">    1</span>&#160;<span class="preprocessor">#pragma once</span></div><div class="line"><a name="l00002"></a><span class="lineno">    2</span>&#160;</div><div class="line"><a name="l00003"></a><span class="lineno">    3</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="cuda__task_8hpp.html">cuda_task.hpp</a>&quot;</span></div><div class="line"><a name="l00004"></a><span class="lineno">    4</span>&#160;<span class="preprocessor">#include &quot;cuda_algorithm/cuda_for_each.hpp&quot;</span></div><div class="line"><a name="l00005"></a><span class="lineno">    5</span>&#160;<span class="preprocessor">#include &quot;cuda_algorithm/cuda_transform.hpp&quot;</span></div><div class="line"><a name="l00006"></a><span class="lineno">    6</span>&#160;</div><div class="line"><a name="l00011"></a><span class="lineno">   11</span>&#160;<span class="keyword">namespace </span><a class="code" href="namespacetf.html">tf</a> {</div><div class="line"><a name="l00012"></a><span class="lineno">   12</span>&#160;</div><div class="line"><a name="l00016"></a><span class="lineno">   16</span>&#160;constexpr <span class="keywordtype">size_t</span> cuda_default_threads_per_block(<span class="keywordtype">size_t</span> N) {</div><div class="line"><a name="l00017"></a><span class="lineno">   17</span>&#160;  <span class="keywordflow">return</span> N &gt;= 256 ? 256 : 128;</div><div class="line"><a name="l00018"></a><span class="lineno">   18</span>&#160;}</div><div class="line"><a name="l00019"></a><span class="lineno">   19</span>&#160;</div><div class="line"><a name="l00020"></a><span class="lineno">   20</span>&#160;<span class="comment">// ----------------------------------------------------------------------------</span></div><div class="line"><a name="l00021"></a><span class="lineno">   21</span>&#160;<span class="comment">// class definition: cudaFlowCapturerBase</span></div><div class="line"><a name="l00022"></a><span class="lineno">   22</span>&#160;<span class="comment">// ----------------------------------------------------------------------------</span></div><div class="line"><a name="l00023"></a><span class="lineno">   23</span>&#160;</div><div class="line"><a name="l00029"></a><span class="lineno"><a class="line" href="classtf_1_1cudaFlowCapturerBase.html">   29</a></span>&#160;<span class="keyword">class </span><a class="code" href="classtf_1_1cudaFlowCapturerBase.html">cudaFlowCapturerBase</a> {</div><div class="line"><a name="l00030"></a><span class="lineno">   30</span>&#160;</div><div class="line"><a name="l00031"></a><span class="lineno">   31</span>&#160;  <span class="keyword">friend</span> <span class="keyword">class </span><a class="code" href="classtf_1_1cudaFlowCapturer.html">cudaFlowCapturer</a>;</div><div class="line"><a name="l00032"></a><span class="lineno">   32</span>&#160;</div><div class="line"><a name="l00033"></a><span class="lineno">   33</span>&#160;  <span class="keyword">public</span>:</div><div class="line"><a name="l00034"></a><span class="lineno">   34</span>&#160;</div><div class="line"><a name="l00038"></a><span class="lineno">   38</span>&#160;    <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a43739b92f51358bd144270664c6fdf8e">cudaFlowCapturerBase</a>() = <span class="keywordflow">default</span>;</div><div class="line"><a name="l00039"></a><span class="lineno">   39</span>&#160;</div><div class="line"><a name="l00043"></a><span class="lineno">   43</span>&#160;    <span class="keyword">virtual</span> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a09d58616acdfe01dfe09271f8b3587ff">~cudaFlowCapturerBase</a>() = <span class="keywordflow">default</span>;</div><div class="line"><a name="l00044"></a><span class="lineno">   44</span>&#160;</div><div class="line"><a name="l00045"></a><span class="lineno">   45</span>&#160;    </div><div class="line"><a name="l00055"></a><span class="lineno">   55</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> C&gt;</div><div class="line"><a name="l00056"></a><span class="lineno">   56</span>&#160;    <a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>(C&amp;&amp; callable);</div><div class="line"><a name="l00057"></a><span class="lineno">   57</span>&#160;    </div><div class="line"><a name="l00068"></a><span class="lineno">   68</span>&#160;    <a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a38adf66dfcc0829708db653d153a83e2">memcpy</a>(<span class="keywordtype">void</span>* dst, <span class="keyword">const</span> <span class="keywordtype">void</span>* src, <span class="keywordtype">size_t</span> count);</div><div class="line"><a name="l00069"></a><span class="lineno">   69</span>&#160;</div><div class="line"><a name="l00084"></a><span class="lineno">   84</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, </div><div class="line"><a name="l00085"></a><span class="lineno">   85</span>&#160;      std::enable_if_t&lt;!std::is_same_v&lt;T, void&gt;, <span class="keywordtype">void</span>&gt;* = <span class="keyword">nullptr</span></div><div class="line"><a name="l00086"></a><span class="lineno">   86</span>&#160;    &gt;</div><div class="line"><a name="l00087"></a><span class="lineno">   87</span>&#160;    <a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a9fd8983c782f841bba6b10de39becdf6">copy</a>(T* tgt, <span class="keyword">const</span> T* src, <span class="keywordtype">size_t</span> num);</div><div class="line"><a name="l00088"></a><span class="lineno">   88</span>&#160;</div><div class="line"><a name="l00100"></a><span class="lineno">  100</span>&#160;    <a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a8e3071171c0875c93dcc077a2e0a435a">memset</a>(<span class="keywordtype">void</span>* devPtr, <span class="keywordtype">int</span> value, <span class="keywordtype">size_t</span> count);</div><div class="line"><a name="l00101"></a><span class="lineno">  101</span>&#160;</div><div class="line"><a name="l00116"></a><span class="lineno">  116</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> F, <span class="keyword">typename</span>... ArgsT&gt;</div><div class="line"><a name="l00117"></a><span class="lineno">  117</span>&#160;    <a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#ad836d32cc2e9532ce57fe3ad6cc67d5d">kernel</a>(dim3 g, dim3 b, <span class="keywordtype">size_t</span> s, F&amp;&amp; f, ArgsT&amp;&amp;... args);</div><div class="line"><a name="l00118"></a><span class="lineno">  118</span>&#160;    </div><div class="line"><a name="l00119"></a><span class="lineno">  119</span>&#160;    <span class="comment">// ------------------------------------------------------------------------</span></div><div class="line"><a name="l00120"></a><span class="lineno">  120</span>&#160;    <span class="comment">// generic algorithms</span></div><div class="line"><a name="l00121"></a><span class="lineno">  121</span>&#160;    <span class="comment">// ------------------------------------------------------------------------</span></div><div class="line"><a name="l00122"></a><span class="lineno">  122</span>&#160;    </div><div class="line"><a name="l00130"></a><span class="lineno">  130</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> C&gt;</div><div class="line"><a name="l00131"></a><span class="lineno">  131</span>&#160;    <a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#ac25236625772a0a2c66831517b9ec288">single_task</a>(C&amp;&amp; callable);</div><div class="line"><a name="l00132"></a><span class="lineno">  132</span>&#160;    </div><div class="line"><a name="l00154"></a><span class="lineno">  154</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> I, <span class="keyword">typename</span> C&gt;</div><div class="line"><a name="l00155"></a><span class="lineno">  155</span>&#160;    <a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a583e4df1139cf6f20b6e79dfde977a51">for_each</a>(I first, I last, C&amp;&amp; callable);</div><div class="line"><a name="l00156"></a><span class="lineno">  156</span>&#160;</div><div class="line"><a name="l00185"></a><span class="lineno">  185</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> I, <span class="keyword">typename</span> C&gt;</div><div class="line"><a name="l00186"></a><span class="lineno">  186</span>&#160;    <a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a0320596ee37d68cfce5746027ebc97de">for_each_index</a>(I first, I last, I step, C&amp;&amp; callable);</div><div class="line"><a name="l00187"></a><span class="lineno">  187</span>&#160;  </div><div class="line"><a name="l00211"></a><span class="lineno">  211</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> I, <span class="keyword">typename</span> C, <span class="keyword">typename</span>... S&gt;</div><div class="line"><a name="l00212"></a><span class="lineno">  212</span>&#160;    <a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a44fb0c626c46de1bb95369e33194f5c7">transform</a>(I first, I last, C&amp;&amp; callable, S... srcs);</div><div class="line"><a name="l00213"></a><span class="lineno">  213</span>&#160;</div><div class="line"><a name="l00214"></a><span class="lineno">  214</span>&#160;  <span class="keyword">private</span>:</div><div class="line"><a name="l00215"></a><span class="lineno">  215</span>&#160;</div><div class="line"><a name="l00216"></a><span class="lineno">  216</span>&#160;    cudaGraph* _graph {<span class="keyword">nullptr</span>};</div><div class="line"><a name="l00217"></a><span class="lineno">  217</span>&#160;</div><div class="line"><a name="l00218"></a><span class="lineno">  218</span>&#160;    <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a43739b92f51358bd144270664c6fdf8e">cudaFlowCapturerBase</a>(cudaGraph&amp;);</div><div class="line"><a name="l00219"></a><span class="lineno">  219</span>&#160;};</div><div class="line"><a name="l00220"></a><span class="lineno">  220</span>&#160;</div><div class="line"><a name="l00221"></a><span class="lineno">  221</span>&#160;<span class="comment">// Constructor</span></div><div class="line"><a name="l00222"></a><span class="lineno">  222</span>&#160;<span class="keyword">inline</span> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a43739b92f51358bd144270664c6fdf8e">cudaFlowCapturerBase::cudaFlowCapturerBase</a>(cudaGraph&amp; g) :</div><div class="line"><a name="l00223"></a><span class="lineno">  223</span>&#160;  _graph {&amp;g} {</div><div class="line"><a name="l00224"></a><span class="lineno">  224</span>&#160;}</div><div class="line"><a name="l00225"></a><span class="lineno">  225</span>&#160;</div><div class="line"><a name="l00226"></a><span class="lineno">  226</span>&#160;<span class="comment">// Function: capture</span></div><div class="line"><a name="l00227"></a><span class="lineno">  227</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> C&gt;</div><div class="line"><a name="l00228"></a><span class="lineno"><a class="line" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">  228</a></span>&#160;<a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">cudaFlowCapturerBase::on</a>(C&amp;&amp; callable) {</div><div class="line"><a name="l00229"></a><span class="lineno">  229</span>&#160;  <span class="keyword">auto</span> node = _graph-&gt;emplace_back(*_graph,</div><div class="line"><a name="l00230"></a><span class="lineno">  230</span>&#160;    std::in_place_type_t&lt;cudaNode::Capture&gt;{}, std::forward&lt;C&gt;(callable)</div><div class="line"><a name="l00231"></a><span class="lineno">  231</span>&#160;  );</div><div class="line"><a name="l00232"></a><span class="lineno">  232</span>&#160;  <span class="keywordflow">return</span> <a class="code" href="classtf_1_1cudaTask.html">cudaTask</a>(node);</div><div class="line"><a name="l00233"></a><span class="lineno">  233</span>&#160;}</div><div class="line"><a name="l00234"></a><span class="lineno">  234</span>&#160;</div><div class="line"><a name="l00235"></a><span class="lineno">  235</span>&#160;<span class="comment">// Function: memcpy</span></div><div class="line"><a name="l00236"></a><span class="lineno"><a class="line" href="classtf_1_1cudaFlowCapturerBase.html#a38adf66dfcc0829708db653d153a83e2">  236</a></span>&#160;<span class="keyword">inline</span> <a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a38adf66dfcc0829708db653d153a83e2">cudaFlowCapturerBase::memcpy</a>(</div><div class="line"><a name="l00237"></a><span class="lineno">  237</span>&#160;  <span class="keywordtype">void</span>* dst, <span class="keyword">const</span> <span class="keywordtype">void</span>* src, <span class="keywordtype">size_t</span> count</div><div class="line"><a name="l00238"></a><span class="lineno">  238</span>&#160;) {</div><div class="line"><a name="l00239"></a><span class="lineno">  239</span>&#160;  <span class="keywordflow">return</span> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([dst, src, count] (cudaStream_t stream) <span class="keyword">mutable</span> {</div><div class="line"><a name="l00240"></a><span class="lineno">  240</span>&#160;    TF_CHECK_CUDA(</div><div class="line"><a name="l00241"></a><span class="lineno">  241</span>&#160;      cudaMemcpyAsync(dst, src, count, cudaMemcpyDefault, stream),</div><div class="line"><a name="l00242"></a><span class="lineno">  242</span>&#160;      <span class="stringliteral">&quot;failed to capture memcpy&quot;</span></div><div class="line"><a name="l00243"></a><span class="lineno">  243</span>&#160;    );</div><div class="line"><a name="l00244"></a><span class="lineno">  244</span>&#160;  });</div><div class="line"><a name="l00245"></a><span class="lineno">  245</span>&#160;}</div><div class="line"><a name="l00246"></a><span class="lineno">  246</span>&#160;    </div><div class="line"><a name="l00247"></a><span class="lineno">  247</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, std::enable_if_t&lt;!std::is_same_v&lt;T, <span class="keywordtype">void</span>&gt;, <span class="keywordtype">void</span>&gt;*&gt;</div><div class="line"><a name="l00248"></a><span class="lineno"><a class="line" href="classtf_1_1cudaFlowCapturerBase.html#a9fd8983c782f841bba6b10de39becdf6">  248</a></span>&#160;<a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a9fd8983c782f841bba6b10de39becdf6">cudaFlowCapturerBase::copy</a>(T* tgt, <span class="keyword">const</span> T* src, <span class="keywordtype">size_t</span> num) {</div><div class="line"><a name="l00249"></a><span class="lineno">  249</span>&#160;  <span class="keywordflow">return</span> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([tgt, src, num] (cudaStream_t stream) <span class="keyword">mutable</span> {</div><div class="line"><a name="l00250"></a><span class="lineno">  250</span>&#160;    TF_CHECK_CUDA(</div><div class="line"><a name="l00251"></a><span class="lineno">  251</span>&#160;      cudaMemcpyAsync(tgt, src, <span class="keyword">sizeof</span>(T)*num, cudaMemcpyDefault, stream),</div><div class="line"><a name="l00252"></a><span class="lineno">  252</span>&#160;      <span class="stringliteral">&quot;failed to capture copy&quot;</span></div><div class="line"><a name="l00253"></a><span class="lineno">  253</span>&#160;    );</div><div class="line"><a name="l00254"></a><span class="lineno">  254</span>&#160;  });</div><div class="line"><a name="l00255"></a><span class="lineno">  255</span>&#160;}</div><div class="line"><a name="l00256"></a><span class="lineno">  256</span>&#160;</div><div class="line"><a name="l00257"></a><span class="lineno">  257</span>&#160;<span class="comment">// Function: memset</span></div><div class="line"><a name="l00258"></a><span class="lineno"><a class="line" href="classtf_1_1cudaFlowCapturerBase.html#a8e3071171c0875c93dcc077a2e0a435a">  258</a></span>&#160;<span class="keyword">inline</span> <a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a8e3071171c0875c93dcc077a2e0a435a">cudaFlowCapturerBase::memset</a>(<span class="keywordtype">void</span>* ptr, <span class="keywordtype">int</span> v, <span class="keywordtype">size_t</span> n) {</div><div class="line"><a name="l00259"></a><span class="lineno">  259</span>&#160;  <span class="keywordflow">return</span> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([ptr, v, n] (cudaStream_t stream) <span class="keyword">mutable</span> {</div><div class="line"><a name="l00260"></a><span class="lineno">  260</span>&#160;    TF_CHECK_CUDA(</div><div class="line"><a name="l00261"></a><span class="lineno">  261</span>&#160;      cudaMemsetAsync(ptr, v, n, stream), <span class="stringliteral">&quot;failed to capture memset&quot;</span></div><div class="line"><a name="l00262"></a><span class="lineno">  262</span>&#160;    );</div><div class="line"><a name="l00263"></a><span class="lineno">  263</span>&#160;  });</div><div class="line"><a name="l00264"></a><span class="lineno">  264</span>&#160;}</div><div class="line"><a name="l00265"></a><span class="lineno">  265</span>&#160;    </div><div class="line"><a name="l00266"></a><span class="lineno">  266</span>&#160;<span class="comment">// Function: kernel</span></div><div class="line"><a name="l00267"></a><span class="lineno">  267</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> F, <span class="keyword">typename</span>... ArgsT&gt;</div><div class="line"><a name="l00268"></a><span class="lineno"><a class="line" href="classtf_1_1cudaFlowCapturerBase.html#ad836d32cc2e9532ce57fe3ad6cc67d5d">  268</a></span>&#160;<a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#ad836d32cc2e9532ce57fe3ad6cc67d5d">cudaFlowCapturerBase::kernel</a>(</div><div class="line"><a name="l00269"></a><span class="lineno">  269</span>&#160;  dim3 g, dim3 b, <span class="keywordtype">size_t</span> shm, F&amp;&amp; f, ArgsT&amp;&amp;... args</div><div class="line"><a name="l00270"></a><span class="lineno">  270</span>&#160;) {</div><div class="line"><a name="l00271"></a><span class="lineno">  271</span>&#160;  <span class="keywordflow">return</span> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([g, b, shm, f, args...] (cudaStream_t stream) <span class="keyword">mutable</span> {</div><div class="line"><a name="l00272"></a><span class="lineno">  272</span>&#160;    f&lt;&lt;&lt;g, b, shm, stream&gt;&gt;&gt;(args...);</div><div class="line"><a name="l00273"></a><span class="lineno">  273</span>&#160;  });</div><div class="line"><a name="l00274"></a><span class="lineno">  274</span>&#160;}</div><div class="line"><a name="l00275"></a><span class="lineno">  275</span>&#160;</div><div class="line"><a name="l00276"></a><span class="lineno">  276</span>&#160;<span class="comment">// Function: single_task</span></div><div class="line"><a name="l00277"></a><span class="lineno">  277</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> C&gt;</div><div class="line"><a name="l00278"></a><span class="lineno"><a class="line" href="classtf_1_1cudaFlowCapturerBase.html#ac25236625772a0a2c66831517b9ec288">  278</a></span>&#160;<a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#ac25236625772a0a2c66831517b9ec288">cudaFlowCapturerBase::single_task</a>(C&amp;&amp; callable) {</div><div class="line"><a name="l00279"></a><span class="lineno">  279</span>&#160;  <span class="keywordflow">return</span> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([c=std::forward&lt;C&gt;(callable)] (cudaStream_t stream) <span class="keyword">mutable</span> {</div><div class="line"><a name="l00280"></a><span class="lineno">  280</span>&#160;    cuda_single_task&lt;C&gt;&lt;&lt;&lt;1, 1, 0, stream&gt;&gt;&gt;(c);</div><div class="line"><a name="l00281"></a><span class="lineno">  281</span>&#160;  });</div><div class="line"><a name="l00282"></a><span class="lineno">  282</span>&#160;}</div><div class="line"><a name="l00283"></a><span class="lineno">  283</span>&#160;</div><div class="line"><a name="l00284"></a><span class="lineno">  284</span>&#160;<span class="comment">// Function: for_each</span></div><div class="line"><a name="l00285"></a><span class="lineno">  285</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> I, <span class="keyword">typename</span> C&gt;</div><div class="line"><a name="l00286"></a><span class="lineno"><a class="line" href="classtf_1_1cudaFlowCapturerBase.html#a583e4df1139cf6f20b6e79dfde977a51">  286</a></span>&#160;<a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a583e4df1139cf6f20b6e79dfde977a51">cudaFlowCapturerBase::for_each</a>(I first, I last, C&amp;&amp; c) {</div><div class="line"><a name="l00287"></a><span class="lineno">  287</span>&#160;  <span class="keywordflow">return</span> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([first, last, c=std::forward&lt;C&gt;(c)](cudaStream_t stream) <span class="keyword">mutable</span> {</div><div class="line"><a name="l00288"></a><span class="lineno">  288</span>&#160;    <span class="keywordtype">size_t</span> N = std::distance(first, last);</div><div class="line"><a name="l00289"></a><span class="lineno">  289</span>&#160;    <span class="keywordtype">size_t</span> B = cuda_default_threads_per_block(N);</div><div class="line"><a name="l00290"></a><span class="lineno">  290</span>&#160;    cuda_for_each&lt;I, C&gt;&lt;&lt;&lt;(N+B-1)/B, B, 0, stream&gt;&gt;&gt;(first, N, c);</div><div class="line"><a name="l00291"></a><span class="lineno">  291</span>&#160;  });</div><div class="line"><a name="l00292"></a><span class="lineno">  292</span>&#160;}</div><div class="line"><a name="l00293"></a><span class="lineno">  293</span>&#160;</div><div class="line"><a name="l00294"></a><span class="lineno">  294</span>&#160;<span class="comment">// Function: for_each_index</span></div><div class="line"><a name="l00295"></a><span class="lineno">  295</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> I, <span class="keyword">typename</span> C&gt;</div><div class="line"><a name="l00296"></a><span class="lineno"><a class="line" href="classtf_1_1cudaFlowCapturerBase.html#a0320596ee37d68cfce5746027ebc97de">  296</a></span>&#160;<a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a0320596ee37d68cfce5746027ebc97de">cudaFlowCapturerBase::for_each_index</a>(I beg, I end, I inc, C&amp;&amp; c) {</div><div class="line"><a name="l00297"></a><span class="lineno">  297</span>&#160;      </div><div class="line"><a name="l00298"></a><span class="lineno">  298</span>&#160;  <span class="keywordflow">if</span>(is_range_invalid(beg, end, inc)) {</div><div class="line"><a name="l00299"></a><span class="lineno">  299</span>&#160;    TF_THROW(<span class="stringliteral">&quot;invalid range [&quot;</span>, beg, <span class="stringliteral">&quot;, &quot;</span>, end, <span class="stringliteral">&quot;) with inc size &quot;</span>, inc);</div><div class="line"><a name="l00300"></a><span class="lineno">  300</span>&#160;  }</div><div class="line"><a name="l00301"></a><span class="lineno">  301</span>&#160;  </div><div class="line"><a name="l00302"></a><span class="lineno">  302</span>&#160;  <span class="keywordflow">return</span> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([beg, end, inc, c=std::forward&lt;C&gt;(c)] (cudaStream_t stream) <span class="keyword">mutable</span> {</div><div class="line"><a name="l00303"></a><span class="lineno">  303</span>&#160;    <span class="comment">// TODO: special case when N is 0?</span></div><div class="line"><a name="l00304"></a><span class="lineno">  304</span>&#160;    <span class="keywordtype">size_t</span> N = distance(beg, end, inc);</div><div class="line"><a name="l00305"></a><span class="lineno">  305</span>&#160;    <span class="keywordtype">size_t</span> B = cuda_default_threads_per_block(N);</div><div class="line"><a name="l00306"></a><span class="lineno">  306</span>&#160;    cuda_for_each_index&lt;I, C&gt;&lt;&lt;&lt;(N+B-1)/B, B, 0, stream&gt;&gt;&gt;(beg, inc, N, c);</div><div class="line"><a name="l00307"></a><span class="lineno">  307</span>&#160;  });</div><div class="line"><a name="l00308"></a><span class="lineno">  308</span>&#160;}</div><div class="line"><a name="l00309"></a><span class="lineno">  309</span>&#160;</div><div class="line"><a name="l00310"></a><span class="lineno">  310</span>&#160;<span class="comment">// Function: transform</span></div><div class="line"><a name="l00311"></a><span class="lineno">  311</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> I, <span class="keyword">typename</span> C, <span class="keyword">typename</span>... S&gt;</div><div class="line"><a name="l00312"></a><span class="lineno"><a class="line" href="classtf_1_1cudaFlowCapturerBase.html#a44fb0c626c46de1bb95369e33194f5c7">  312</a></span>&#160;<a class="code" href="classtf_1_1cudaTask.html">cudaTask</a> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#a44fb0c626c46de1bb95369e33194f5c7">cudaFlowCapturerBase::transform</a>(I first, I last, C&amp;&amp; c, S... srcs) {</div><div class="line"><a name="l00313"></a><span class="lineno">  313</span>&#160;  <span class="keywordflow">return</span> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">on</a>([first, last, c=std::forward&lt;C&gt;(c), srcs...] (cudaStream_t stream) <span class="keyword">mutable</span> {</div><div class="line"><a name="l00314"></a><span class="lineno">  314</span>&#160;    <span class="comment">// TODO: special case when N is 0?</span></div><div class="line"><a name="l00315"></a><span class="lineno">  315</span>&#160;    size_t N = std::distance(first, last);</div><div class="line"><a name="l00316"></a><span class="lineno">  316</span>&#160;    size_t B = cuda_default_threads_per_block(N);</div><div class="line"><a name="l00317"></a><span class="lineno">  317</span>&#160;    cuda_transform&lt;I, C, S...&gt;&lt;&lt;&lt;(N+B-1)/B, B, 0, stream&gt;&gt;&gt;(first, N, c, srcs...);</div><div class="line"><a name="l00318"></a><span class="lineno">  318</span>&#160;  });</div><div class="line"><a name="l00319"></a><span class="lineno">  319</span>&#160;}</div><div class="line"><a name="l00320"></a><span class="lineno">  320</span>&#160;</div><div class="line"><a name="l00321"></a><span class="lineno">  321</span>&#160;<span class="comment">// ----------------------------------------------------------------------------</span></div><div class="line"><a name="l00322"></a><span class="lineno">  322</span>&#160;<span class="comment">// class definition: cudaFlowCapturer</span></div><div class="line"><a name="l00323"></a><span class="lineno">  323</span>&#160;<span class="comment">// ----------------------------------------------------------------------------</span></div><div class="line"><a name="l00324"></a><span class="lineno">  324</span>&#160;</div><div class="line"><a name="l00359"></a><span class="lineno"><a class="line" href="classtf_1_1cudaFlowCapturer.html">  359</a></span>&#160;<span class="keyword">class </span><a class="code" href="classtf_1_1cudaFlowCapturer.html">cudaFlowCapturer</a> : <span class="keyword">public</span> <a class="code" href="classtf_1_1cudaFlowCapturerBase.html">cudaFlowCapturerBase</a> {</div><div class="line"><a name="l00360"></a><span class="lineno">  360</span>&#160;</div><div class="line"><a name="l00361"></a><span class="lineno">  361</span>&#160;  <span class="keyword">friend</span> <span class="keyword">class </span><a class="code" href="classtf_1_1cudaFlow.html">cudaFlow</a>;</div><div class="line"><a name="l00362"></a><span class="lineno">  362</span>&#160;  <span class="keyword">friend</span> <span class="keyword">class </span><a class="code" href="classtf_1_1Executor.html">Executor</a>;</div><div class="line"><a name="l00363"></a><span class="lineno">  363</span>&#160;</div><div class="line"><a name="l00364"></a><span class="lineno">  364</span>&#160;  <span class="keyword">public</span>:</div><div class="line"><a name="l00365"></a><span class="lineno">  365</span>&#160;    </div><div class="line"><a name="l00369"></a><span class="lineno">  369</span>&#160;    <span class="keywordtype">bool</span> <a class="code" href="classtf_1_1cudaFlowCapturer.html#a3413a20a7c8229365e1ee9fb5af4af1e">empty</a>() <span class="keyword">const</span>;</div><div class="line"><a name="l00370"></a><span class="lineno">  370</span>&#160;    </div><div class="line"><a name="l00383"></a><span class="lineno">  383</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, <span class="keyword">typename</span>... ArgsT&gt;</div><div class="line"><a name="l00384"></a><span class="lineno">  384</span>&#160;    T* <a class="code" href="classtf_1_1cudaFlowCapturer.html#a49dd9727aeb5bccd7d826c6b11a6882d">make_capturer</a>(ArgsT&amp;&amp;... args);</div><div class="line"><a name="l00385"></a><span class="lineno">  385</span>&#160;</div><div class="line"><a name="l00386"></a><span class="lineno">  386</span>&#160;  <span class="keyword">private</span>:</div><div class="line"><a name="l00387"></a><span class="lineno">  387</span>&#160;    </div><div class="line"><a name="l00388"></a><span class="lineno">  388</span>&#160;    <a class="codeRef" doxygen="/Users/twhuang/PhD/Code/taskflow/doxygen/cppreference-doxygen-web.tag.xml:http://en.cppreference.com/w/" href="http://en.cppreference.com/w/cpp/container/vector.html">std::vector&lt;std::unique_ptr&lt;cudaFlowCapturerBase&gt;</a>&gt; _capturers;</div><div class="line"><a name="l00389"></a><span class="lineno">  389</span>&#160;</div><div class="line"><a name="l00390"></a><span class="lineno">  390</span>&#160;    <a class="code" href="classtf_1_1cudaFlowCapturer.html">cudaFlowCapturer</a>(cudaGraph&amp;);</div><div class="line"><a name="l00391"></a><span class="lineno">  391</span>&#160;    </div><div class="line"><a name="l00392"></a><span class="lineno">  392</span>&#160;    cudaGraph_t _capture();</div><div class="line"><a name="l00393"></a><span class="lineno">  393</span>&#160;    </div><div class="line"><a name="l00394"></a><span class="lineno">  394</span>&#160;    cudaGraphExec_t _executable {<span class="keyword">nullptr</span>};</div><div class="line"><a name="l00395"></a><span class="lineno">  395</span>&#160;    </div><div class="line"><a name="l00396"></a><span class="lineno">  396</span>&#160;    <span class="comment">//void _create_executable();</span></div><div class="line"><a name="l00397"></a><span class="lineno">  397</span>&#160;    <span class="comment">//void _destroy_executable();</span></div><div class="line"><a name="l00398"></a><span class="lineno">  398</span>&#160;};</div><div class="line"><a name="l00399"></a><span class="lineno">  399</span>&#160;</div><div class="line"><a name="l00400"></a><span class="lineno">  400</span>&#160;<span class="comment">// constructor</span></div><div class="line"><a name="l00401"></a><span class="lineno">  401</span>&#160;<span class="keyword">inline</span> cudaFlowCapturer::cudaFlowCapturer(cudaGraph&amp; g) :</div><div class="line"><a name="l00402"></a><span class="lineno">  402</span>&#160;  <a class="code" href="classtf_1_1cudaFlowCapturerBase.html">cudaFlowCapturerBase</a>{g} {</div><div class="line"><a name="l00403"></a><span class="lineno">  403</span>&#160;}</div><div class="line"><a name="l00404"></a><span class="lineno">  404</span>&#160;</div><div class="line"><a name="l00405"></a><span class="lineno">  405</span>&#160;<span class="comment">// Function: empty</span></div><div class="line"><a name="l00406"></a><span class="lineno"><a class="line" href="classtf_1_1cudaFlowCapturer.html#a3413a20a7c8229365e1ee9fb5af4af1e">  406</a></span>&#160;<span class="keyword">inline</span> <span class="keywordtype">bool</span> <a class="code" href="classtf_1_1cudaFlowCapturer.html#a3413a20a7c8229365e1ee9fb5af4af1e">cudaFlowCapturer::empty</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00407"></a><span class="lineno">  407</span>&#160;  <span class="keywordflow">return</span> _graph-&gt;empty();</div><div class="line"><a name="l00408"></a><span class="lineno">  408</span>&#160;}</div><div class="line"><a name="l00409"></a><span class="lineno">  409</span>&#160;</div><div class="line"><a name="l00411"></a><span class="lineno">  411</span>&#160;<span class="comment">//inline void cudaFlowCapturer::_create_executable() {</span></div><div class="line"><a name="l00412"></a><span class="lineno">  412</span>&#160;<span class="comment">//  assert(_executable == nullptr);</span></div><div class="line"><a name="l00413"></a><span class="lineno">  413</span>&#160;<span class="comment">//  TF_CHECK_CUDA(</span></div><div class="line"><a name="l00414"></a><span class="lineno">  414</span>&#160;<span class="comment">//    cudaGraphInstantiate(</span></div><div class="line"><a name="l00415"></a><span class="lineno">  415</span>&#160;<span class="comment">//      &amp;_executable, _graph-&gt;_native_handle, nullptr, nullptr, 0</span></div><div class="line"><a name="l00416"></a><span class="lineno">  416</span>&#160;<span class="comment">//    ),</span></div><div class="line"><a name="l00417"></a><span class="lineno">  417</span>&#160;<span class="comment">//    &quot;failed to create an executable for captured graph&quot;</span></div><div class="line"><a name="l00418"></a><span class="lineno">  418</span>&#160;<span class="comment">//  );</span></div><div class="line"><a name="l00419"></a><span class="lineno">  419</span>&#160;<span class="comment">//}</span></div><div class="line"><a name="l00420"></a><span class="lineno">  420</span>&#160;<span class="comment">//</span></div><div class="line"><a name="l00422"></a><span class="lineno">  422</span>&#160;<span class="comment"></span><span class="comment">//inline void cudaFlowCapturer::_destroy_executable() {</span></div><div class="line"><a name="l00423"></a><span class="lineno">  423</span>&#160;<span class="comment">//  assert(_executable != nullptr);</span></div><div class="line"><a name="l00424"></a><span class="lineno">  424</span>&#160;<span class="comment">//  TF_CHECK_CUDA(</span></div><div class="line"><a name="l00425"></a><span class="lineno">  425</span>&#160;<span class="comment">//    cudaGraphExecDestroy(_executable), &quot;failed to destroy executable graph&quot;</span></div><div class="line"><a name="l00426"></a><span class="lineno">  426</span>&#160;<span class="comment">//  );</span></div><div class="line"><a name="l00427"></a><span class="lineno">  427</span>&#160;<span class="comment">//  _executable = nullptr;</span></div><div class="line"><a name="l00428"></a><span class="lineno">  428</span>&#160;<span class="comment">//}</span></div><div class="line"><a name="l00429"></a><span class="lineno">  429</span>&#160;</div><div class="line"><a name="l00430"></a><span class="lineno">  430</span>&#160;<span class="comment">// Function: make_capturer</span></div><div class="line"><a name="l00431"></a><span class="lineno">  431</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, <span class="keyword">typename</span>... ArgsT&gt;</div><div class="line"><a name="l00432"></a><span class="lineno"><a class="line" href="classtf_1_1cudaFlowCapturer.html#a49dd9727aeb5bccd7d826c6b11a6882d">  432</a></span>&#160;T* <a class="code" href="classtf_1_1cudaFlowCapturer.html#a49dd9727aeb5bccd7d826c6b11a6882d">cudaFlowCapturer::make_capturer</a>(ArgsT&amp;&amp;... args) {</div><div class="line"><a name="l00433"></a><span class="lineno">  433</span>&#160;</div><div class="line"><a name="l00434"></a><span class="lineno">  434</span>&#160;  static_assert(std::is_base_of_v&lt;cudaFlowCapturerBase, T&gt;);</div><div class="line"><a name="l00435"></a><span class="lineno">  435</span>&#160;</div><div class="line"><a name="l00436"></a><span class="lineno">  436</span>&#160;  <span class="keyword">auto</span> ptr = std::make_unique&lt;T&gt;(std::forward&lt;ArgsT&gt;(args)...);</div><div class="line"><a name="l00437"></a><span class="lineno">  437</span>&#160;  ptr-&gt;_graph = this-&gt;_graph;</div><div class="line"><a name="l00438"></a><span class="lineno">  438</span>&#160;  <span class="keyword">auto</span> raw = ptr.get();</div><div class="line"><a name="l00439"></a><span class="lineno">  439</span>&#160;  _capturers.push_back(std::move(ptr));</div><div class="line"><a name="l00440"></a><span class="lineno">  440</span>&#160;  <span class="keywordflow">return</span> raw;</div><div class="line"><a name="l00441"></a><span class="lineno">  441</span>&#160;}</div><div class="line"><a name="l00442"></a><span class="lineno">  442</span>&#160;</div><div class="line"><a name="l00443"></a><span class="lineno">  443</span>&#160;<span class="comment">// Function: _capture</span></div><div class="line"><a name="l00444"></a><span class="lineno">  444</span>&#160;<span class="keyword">inline</span> cudaGraph_t cudaFlowCapturer::_capture() {</div><div class="line"><a name="l00445"></a><span class="lineno">  445</span>&#160;</div><div class="line"><a name="l00446"></a><span class="lineno">  446</span>&#160;  <span class="comment">// acquire per-thread stream and turn it into capture mode</span></div><div class="line"><a name="l00447"></a><span class="lineno">  447</span>&#160;  <span class="comment">// we must use ThreadLocal mode to avoid clashing with CUDA global states</span></div><div class="line"><a name="l00448"></a><span class="lineno">  448</span>&#160;  cudaScopedPerThreadStream stream;</div><div class="line"><a name="l00449"></a><span class="lineno">  449</span>&#160;  </div><div class="line"><a name="l00450"></a><span class="lineno">  450</span>&#160;  TF_CHECK_CUDA(</div><div class="line"><a name="l00451"></a><span class="lineno">  451</span>&#160;    cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal), </div><div class="line"><a name="l00452"></a><span class="lineno">  452</span>&#160;    <span class="stringliteral">&quot;failed to turn stream into per-thread capture mode&quot;</span></div><div class="line"><a name="l00453"></a><span class="lineno">  453</span>&#160;  );</div><div class="line"><a name="l00454"></a><span class="lineno">  454</span>&#160;</div><div class="line"><a name="l00455"></a><span class="lineno">  455</span>&#160;  <span class="comment">// TODO: need an efficient algorithm</span></div><div class="line"><a name="l00456"></a><span class="lineno">  456</span>&#160;  <span class="keyword">auto</span> ordered = _graph-&gt;_toposort();</div><div class="line"><a name="l00457"></a><span class="lineno">  457</span>&#160;  <span class="keywordflow">for</span>(<span class="keyword">auto</span>&amp; node : ordered) {</div><div class="line"><a name="l00458"></a><span class="lineno">  458</span>&#160;    std::get&lt;cudaNode::Capture&gt;(node-&gt;_handle).work(stream);  </div><div class="line"><a name="l00459"></a><span class="lineno">  459</span>&#160;  }</div><div class="line"><a name="l00460"></a><span class="lineno">  460</span>&#160;</div><div class="line"><a name="l00461"></a><span class="lineno">  461</span>&#160;  cudaGraph_t g;</div><div class="line"><a name="l00462"></a><span class="lineno">  462</span>&#160;</div><div class="line"><a name="l00463"></a><span class="lineno">  463</span>&#160;  TF_CHECK_CUDA(</div><div class="line"><a name="l00464"></a><span class="lineno">  464</span>&#160;    cudaStreamEndCapture(stream, &amp;g), <span class="stringliteral">&quot;failed to end capture&quot;</span></div><div class="line"><a name="l00465"></a><span class="lineno">  465</span>&#160;  );</div><div class="line"><a name="l00466"></a><span class="lineno">  466</span>&#160;</div><div class="line"><a name="l00467"></a><span class="lineno">  467</span>&#160;  <span class="comment">//cuda_dump_graph(std::cout, g);</span></div><div class="line"><a name="l00468"></a><span class="lineno">  468</span>&#160;  </div><div class="line"><a name="l00469"></a><span class="lineno">  469</span>&#160;  <span class="keywordflow">return</span> g;</div><div class="line"><a name="l00470"></a><span class="lineno">  470</span>&#160;}</div><div class="line"><a name="l00471"></a><span class="lineno">  471</span>&#160;</div><div class="line"><a name="l00472"></a><span class="lineno">  472</span>&#160;}  <span class="comment">// end of namespace tf -----------------------------------------------------</span></div><div class="line"><a name="l00473"></a><span class="lineno">  473</span>&#160;</div><div class="ttc" id="classtf_1_1cudaFlowCapturerBase_html_a8e3071171c0875c93dcc077a2e0a435a"><div class="ttname"><a href="classtf_1_1cudaFlowCapturerBase.html#a8e3071171c0875c93dcc077a2e0a435a">tf::cudaFlowCapturerBase::memset</a></div><div class="ttdeci">cudaTask memset(void *devPtr, int value, size_t count)</div><div class="ttdoc">initializes or sets GPU memory to the given value byte by byte </div><div class="ttdef"><b>Definition:</b> cuda_capturer.hpp:258</div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturerBase_html_adf651356def71f613c589c29588398c2"><div class="ttname"><a href="classtf_1_1cudaFlowCapturerBase.html#adf651356def71f613c589c29588398c2">tf::cudaFlowCapturerBase::on</a></div><div class="ttdeci">cudaTask on(C &amp;&amp;callable)</div><div class="ttdoc">captures a sequential CUDA operations from the given callable </div><div class="ttdef"><b>Definition:</b> cuda_capturer.hpp:228</div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturerBase_html_a583e4df1139cf6f20b6e79dfde977a51"><div class="ttname"><a href="classtf_1_1cudaFlowCapturerBase.html#a583e4df1139cf6f20b6e79dfde977a51">tf::cudaFlowCapturerBase::for_each</a></div><div class="ttdeci">cudaTask for_each(I first, I last, C &amp;&amp;callable)</div><div class="ttdoc">captures a kernel that applies a callable to each dereferenced element of the data array ...</div><div class="ttdef"><b>Definition:</b> cuda_capturer.hpp:286</div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturerBase_html_ac25236625772a0a2c66831517b9ec288"><div class="ttname"><a href="classtf_1_1cudaFlowCapturerBase.html#ac25236625772a0a2c66831517b9ec288">tf::cudaFlowCapturerBase::single_task</a></div><div class="ttdeci">cudaTask single_task(C &amp;&amp;callable)</div><div class="ttdoc">capturers a kernel to runs the given callable with only one thread </div><div class="ttdef"><b>Definition:</b> cuda_capturer.hpp:278</div></div>
<div class="ttc" id="namespacetf_html"><div class="ttname"><a href="namespacetf.html">tf</a></div><div class="ttdef"><b>Definition:</b> flow_builder.hpp:9</div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturerBase_html_a38adf66dfcc0829708db653d153a83e2"><div class="ttname"><a href="classtf_1_1cudaFlowCapturerBase.html#a38adf66dfcc0829708db653d153a83e2">tf::cudaFlowCapturerBase::memcpy</a></div><div class="ttdeci">cudaTask memcpy(void *dst, const void *src, size_t count)</div><div class="ttdoc">copies data between host and device asynchronously through a stream </div><div class="ttdef"><b>Definition:</b> cuda_capturer.hpp:236</div></div>
<div class="ttc" id="cuda__task_8hpp_html"><div class="ttname"><a href="cuda__task_8hpp.html">cuda_task.hpp</a></div></div>
<div class="ttc" id="classtf_1_1cudaFlow_html"><div class="ttname"><a href="classtf_1_1cudaFlow.html">tf::cudaFlow</a></div><div class="ttdoc">class for building a CUDA task dependency graph </div><div class="ttdef"><b>Definition:</b> cuda_flow.hpp:47</div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturerBase_html_a43739b92f51358bd144270664c6fdf8e"><div class="ttname"><a href="classtf_1_1cudaFlowCapturerBase.html#a43739b92f51358bd144270664c6fdf8e">tf::cudaFlowCapturerBase::cudaFlowCapturerBase</a></div><div class="ttdeci">cudaFlowCapturerBase()=default</div><div class="ttdoc">default constructor </div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturerBase_html_a44fb0c626c46de1bb95369e33194f5c7"><div class="ttname"><a href="classtf_1_1cudaFlowCapturerBase.html#a44fb0c626c46de1bb95369e33194f5c7">tf::cudaFlowCapturerBase::transform</a></div><div class="ttdeci">cudaTask transform(I first, I last, C &amp;&amp;callable, S... srcs)</div><div class="ttdoc">captures a kernel that applies a callable to a source range and stores the result in a target range ...</div><div class="ttdef"><b>Definition:</b> cuda_capturer.hpp:312</div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturerBase_html"><div class="ttname"><a href="classtf_1_1cudaFlowCapturerBase.html">tf::cudaFlowCapturerBase</a></div><div class="ttdoc">base class to construct a CUDA task graph through stream capture </div><div class="ttdef"><b>Definition:</b> cuda_capturer.hpp:29</div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturer_html_a49dd9727aeb5bccd7d826c6b11a6882d"><div class="ttname"><a href="classtf_1_1cudaFlowCapturer.html#a49dd9727aeb5bccd7d826c6b11a6882d">tf::cudaFlowCapturer::make_capturer</a></div><div class="ttdeci">T * make_capturer(ArgsT &amp;&amp;... args)</div><div class="ttdoc">creates a custom capturer derived from tf::cudaFlowCapturerBase </div><div class="ttdef"><b>Definition:</b> cuda_capturer.hpp:432</div></div>
<div class="ttc" id="vector_html"><div class="ttname"><a href="http://en.cppreference.com/w/cpp/container/vector.html">std::vector</a></div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturerBase_html_ad836d32cc2e9532ce57fe3ad6cc67d5d"><div class="ttname"><a href="classtf_1_1cudaFlowCapturerBase.html#ad836d32cc2e9532ce57fe3ad6cc67d5d">tf::cudaFlowCapturerBase::kernel</a></div><div class="ttdeci">cudaTask kernel(dim3 g, dim3 b, size_t s, F &amp;&amp;f, ArgsT &amp;&amp;... args)</div><div class="ttdoc">captures a kernel </div><div class="ttdef"><b>Definition:</b> cuda_capturer.hpp:268</div></div>
<div class="ttc" id="classtf_1_1cudaTask_html"><div class="ttname"><a href="classtf_1_1cudaTask.html">tf::cudaTask</a></div><div class="ttdoc">handle to a node of the internal CUDA graph </div><div class="ttdef"><b>Definition:</b> cuda_task.hpp:53</div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturerBase_html_a0320596ee37d68cfce5746027ebc97de"><div class="ttname"><a href="classtf_1_1cudaFlowCapturerBase.html#a0320596ee37d68cfce5746027ebc97de">tf::cudaFlowCapturerBase::for_each_index</a></div><div class="ttdeci">cudaTask for_each_index(I first, I last, I step, C &amp;&amp;callable)</div><div class="ttdoc">captures a kernel that applies a callable to each index in the range with the step size ...</div><div class="ttdef"><b>Definition:</b> cuda_capturer.hpp:296</div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturerBase_html_a09d58616acdfe01dfe09271f8b3587ff"><div class="ttname"><a href="classtf_1_1cudaFlowCapturerBase.html#a09d58616acdfe01dfe09271f8b3587ff">tf::cudaFlowCapturerBase::~cudaFlowCapturerBase</a></div><div class="ttdeci">virtual ~cudaFlowCapturerBase()=default</div><div class="ttdoc">default virtual destructor </div></div>
<div class="ttc" id="classtf_1_1Executor_html"><div class="ttname"><a href="classtf_1_1Executor.html">tf::Executor</a></div><div class="ttdoc">execution interface for running a taskflow graph </div><div class="ttdef"><b>Definition:</b> executor.hpp:28</div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturerBase_html_a9fd8983c782f841bba6b10de39becdf6"><div class="ttname"><a href="classtf_1_1cudaFlowCapturerBase.html#a9fd8983c782f841bba6b10de39becdf6">tf::cudaFlowCapturerBase::copy</a></div><div class="ttdeci">cudaTask copy(T *tgt, const T *src, size_t num)</div><div class="ttdoc">captures a copy task of typed data </div><div class="ttdef"><b>Definition:</b> cuda_capturer.hpp:248</div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturer_html_a3413a20a7c8229365e1ee9fb5af4af1e"><div class="ttname"><a href="classtf_1_1cudaFlowCapturer.html#a3413a20a7c8229365e1ee9fb5af4af1e">tf::cudaFlowCapturer::empty</a></div><div class="ttdeci">bool empty() const</div><div class="ttdoc">queries the emptiness of the graph </div><div class="ttdef"><b>Definition:</b> cuda_capturer.hpp:406</div></div>
<div class="ttc" id="classtf_1_1cudaFlowCapturer_html"><div class="ttname"><a href="classtf_1_1cudaFlowCapturer.html">tf::cudaFlowCapturer</a></div><div class="ttdoc">class for building a CUDA task dependency graph through stream capture </div><div class="ttdef"><b>Definition:</b> cuda_capturer.hpp:359</div></div>
</div><!-- fragment --></div><!-- contents -->
</div><!-- doc-content -->
<!-- start footer part -->
<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
  <ul>
    <li class="navelem"><a class="el" href="dir_88dad41ea55ca2177e141d32a93e931c.html">taskflow</a></li><li class="navelem"><a class="el" href="dir_638d51f8e6f20ea8c720cc8c006296ba.html">cuda</a></li><li class="navelem"><a class="el" href="cuda__capturer_8hpp.html">cuda_capturer.hpp</a></li>
    <li class="footer">Generated by
    <a href="http://www.doxygen.org/index.html">
    <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.14 </li>
  </ul>
</div>
</body>
</html>
